Repository navigation
feat(cuda): Execute Widen arrays on CUDA - #11303
Conversation
| @@ -63,7 +64,7 @@ __device__ __forceinline__ Output decimal_cast_value(Input value) { | |||
| } | |||
There was a problem hiding this comment.
I am unfamiliar with the GPU code but from my LLM conversations, it seems like we do not know if the optimizer can remove intermediate conversions? i.e. i8->i64->i16?
There was a problem hiding this comment.
In what scenario / why would that happen? But I would think so if you directly chain multiple decimal_cast_value, as the branching is done at compile time.
Merging this PR will not alter performance2 warnings
Comparing |
1741725 to
37897fe
Compare
|
@connortsui20 Having a look now, currently working through the sanitizer error. |
) ## Summary Preserve alignment metadata when uploading host buffers through `ensure_on_device` and its synchronous variant. Uploading a typed buffer as bytes currently reduces its recorded alignment to one byte, even when the CUDA allocation satisfies the original alignment. This causes decimal reconstruction to reject otherwise valid device buffers, as exposed by the Widen tests in #11303. Restore the source alignment with a metadata-only update after the existing upload. No additional data copy, device allocation, or synchronization is introduced for alignment. ## Changes - Preserve source alignment in asynchronous and synchronous uploads. Already-device buffers remain unchanged. - Check alignment against the actual device pointer, including the slice offset. Requests above 256 bytes remain supported when the address satisfies them. - Return an error rather than panic or copy when the device pointer cannot satisfy the requested alignment. Uploads now explicitly fail if the source alignment cannot be preserved. - Reduce slice alignment metadata according to its offset instead of retaining an invalid stronger guarantee. - Add regression coverage for all six decimal widths, full/sliced/empty uploads, both upload paths, already-device identity, and metadata-only alignment updates. --------- Signed-off-by: Alexander Droste <alexander.droste@protonmail.com>
|
@claude review the pr but only code from the latest commit, in particular double check whether all tests are meaningful or should be trimmed down. |
|
Claude finished @0ax1's task in 1m 43s —— View job Review of the latest commit only (
|
Codex Review Summary: ✅ CompletedThis comment shows the latest Codex review activity on this pull request.
ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings. |
|
Codex Review: Didn't find any major issues. Breezy! Reviewed commit: ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
If Codex has suggestions, it will comment; otherwise it will react with 👍. Codex can also answer questions or update the PR. Try commenting "@codex address that feedback". |
0ax1
left a comment
There was a problem hiding this comment.
@connortsui20 PR looks good, but we could trim down the amount of unit tests
Device-resident Widen arrays cannot use the CPU fallback. Restore their logical width on the device so narrow integer and decimal children can execute through CUDA. Progress towards #11284. Signed-off-by: Connor Tsui <connor.tsui20@gmail.com> Signed-off-by: Alexander Droste <alexander.droste@protonmail.com>
3be09d1 to
712a4a3
Compare
Summary
Progress towards #11284 (part 1).
CUDA has no kernel for
vortex.widen, and the CPU fallback only runs when every buffer is on the host, so executing a device-residentWidenfails with "CPU fallback with device-resident buffers is not supported". After CD5, every DecimalByteParts array with a narrow MSP reads from a file with an integerWidenas its MSP, so CD5 can't merge until GPU reads can execute it.Changes
Widenis not fused into the dynamic dispatch path. Dynamic dispatch zero-extends signed narrow BitPacked and FoR children, so fusing a signedWidenthere needs care to sign-extend instead. TheTODOinplan_builder.rstracks that. Until then,WidenExecutorexecutes the child on the device and widens it in a separate kernel.Decimal storage that is wider than the
Widen's logical type is an error instead of a narrowing copy on the device. Narrowing needs either an overflow check or trust in the stored precision, and that choice is the open question in part 2 of #11284.The decimal path reuses the kernels in
decimal_cast.cu, and the Arrow export now shareswiden_decimal_valueswith it, so there is one copy of the widening code. The tests compare every integer and decimal widening pair against the CPUWiden(with nulls, negative values, a FoR child, and the kernel's block boundaries). They only run in the GPU CI.